feat(rocm): wvSplitK skinny GEMM for decode M<=4 — the #487 decode GEMM lever - #506

Reviewed as part of a sweep over the open external PRs. The port itself is careful — the upstream anchor is cited correctly and I verified csrc/rocm/skinny_gemms.cu:351-573 and :1169 are both exact at the pinned oracle, and the arithmetic matches the donor. What needs work is not the kernel body but the dispatch preconditions around it, three of which did not come across.

Flagging up front: this is the same shape as findings on #523 and #509, so it is worth treating as one lesson rather than three separate reviews. In each case the calculation was ported faithfully and the guards vLLM performs before choosing that calculation were not.

1. Out-of-bounds device write on an odd output dimension.

Upstream launches wvSplitK_hf_sml_ only under if ((Kbp_in * N_in <= max_lds_len) && (M_in % _YTILE == 0)) (skinny_gemms.cu:1217). Odd M_in falls through to wvSplitK_hf_. The port carried the LDS-fit half of that predicate across as K * M <= 32768 and dropped the % YTILE half.

The store at rocm_skinny_gemm.hip:122 is unguarded:

if (threadIdx.x == (kThrds - 1)) {
  for (int n = 0; n < N; n++)
    for (int y = 0; y < kYtile; y++) C[m + y + n * M] = __float2bfloat16(sum[n][y]);
}

m is always even and the loop condition is m < M. The read is clamped (min__(y + m, M - 1), line 86); the write is not. With an odd output dim the last active wave has m == N-1, and y == 1 writes C[N + n*N] — for n == M-1 that is two bytes past the end of the output buffer. On a discrete card this is silent memory corruption, not a fault. vt::MatmulBT is a general shared-seam op with no documented parity precondition, and the gate at rocm_matmul_hipblaslt.hip:487 routes every qualifying call, so a GPT-2-family lm_head at vocab 50257 reaches it.

2. No arch guard, and wave32 is hardcoded while the path defaults ON.

Upstream guards twice — on_gfx9() or on_gfx1x() at dispatch (utils.py:174) and #if defined(__HIP__GFX9__) || defined(__HIP__GFX1X__) at compile, with an UNREACHABLE_CODE fallback. It also branches the final reduction on wave width: gfx9 needs ROW_BCAST15/ROW_BCAST31 rather than __shfl_xor(sum, 16) (skinny_gemms.cu:489-496).

The port has neither guard and takes only the wave32 arm, while SkinnyGemmEnabled() defaults true. Built with VLLM_CPP_HIP_ARCHITECTURES=gfx942, every decode MatmulBT with M<=4 silently produces wrong numbers with no error — __shfl_xor(x, 16) does not complete a reduction on a 64-wide wave. rocm_backend.hip:75 already reads props.gcnArchName into caps.gcn_arch, so the information is there and unused. The PR body's "Honest gaps" says gfx1100-only; nothing in the code does.

3. The m > 8 lower bound was dropped (utils.py:181: if m > 8 and 0 < n <= 5). The port gates only on M and never bounds N below. This compounds #1: N == 1 gives m = 0, y = 1 → C[1], out of bounds on the very first wave.

4. The upstream test exists and was not ported — and it contains exactly these cases.

tests/kernels/quantization/test_rocm_skinny_gemms.py::test_rocm_wvsplitk_kernel is present at the pin, and its NKM_FACTORS_WVSPLITK list deliberately includes (4, 4096, 4096 + 1), (4, 16384 * 2, 8192 + 1) and a block commented # Minimum M constraint validation (m >= 8). AGENTS.md asks for the upstream test in the same change with parameters preserved; doing that here would have caught findings 1 and 3 before the PR was opened. It also sweeps n ∈ {1,2,3,4} and both bf16 and fp16.

Two smaller things while you are in there. The shipped test bounds aggregate NMSE where upstream bounds elementwise (atol = eps * sqrt(k), rtol = 1e-2) — at the largest shape one arbitrarily-wrong element contributes ~4.9e-5 against a 5e-4 budget, so roughly ten completely wrong outputs still pass, and the tail defect above touches 1–4. And every K in the test is an exact multiple of the 512 stride, so if (k_ >= K) break never fires and the entire K-tail path is uncovered while real shapes hit it every call.

On the merge conflict: only scripts/env-doc-allowlist.txt, but do not reapply it mechanically — main re-sorted and de-duplicated that whole file (235 → 221 lines) while this branch inserted VT_ROCM_SKINNY three times to match the old three-section layout. Take main's version and insert once, in sorted position after VT_ROCM_HIPBLASLT. (That file being a shared must-write surface is our problem, not yours; it is on the list to fix.)

We have no AMD hardware here, so none of your measurements could be reproduced and I am not disputing them — findings 1–3 are read from the source against the donor at the pin, and are independent of any run. Happy to look again once the guards are back.